#pragma OPENCL EXTENSION cl_khr_3d_image_writes : enable
#pragma OPENCL EXTENSION cl_arm_printf : enable
#pragma OPENCL EXTENSION cl_arm_printf : enable
#define ACCUM_FLT4 float4
#define FLT float
#define FLT2 float2
#define FLT3 float3
#define FLT4 float4
#define TO_FLT4 convert_float4
#define TO_ACCUM_TYPE convert_float4
#define TO_ACCUM_FLT convert_float
#define READ_IMAGE read_imagef
#define WRITE_IMAGE write_imagef
__constant sampler_t smp_edge = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST;
__constant sampler_t smp_none = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_NONE | CLK_FILTER_NEAREST;
__constant sampler_t smp_zero = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP | CLK_FILTER_NEAREST;
#define ACCUM_FLT float
constant ACCUM_FLT Bt[36] = {
        1.0000000000f, 0.0000000000f, -2.5000004768f, -0.0000001192f, 1.0000001192f, 0.0000000000f, 
        0.0000000000f, 0.9428091049f, 1.3333333731f, -0.4714044929f, -0.6666667461f, 0.0000000000f, 
        0.0000000000f, -0.9428089857f, 1.3333334923f, 0.4714045525f, -0.6666667461f, 0.0000000000f, 
        0.0000000000f, -0.1178511307f, -0.0833333358f, 0.2357022613f, 0.1666666865f, 0.0000000000f, 
        0.0000000000f, 0.1178511307f, -0.0833333507f, -0.2357022911f, 0.1666666865f, 0.0000000000f, 
        0.0000000000f, 0.9999998808f, -0.0000000596f, -2.5000000000f, 0.0000000000f, 1.0000000000f, 
};
__kernel void winograd_4x4_to_36_tflite(
__global float4* dst_tensor_buffer,
  __global float4* src_tensor_buffer,
  __read_only image2d_t bt_tex2d,
  int4 shared_int4_0,
  int4 shared_int4_1,
  int4 shared_int4_2) {
  int DST_X = get_global_id(0);
  int DST_Y = get_global_id(1);
  int DST_Z = get_global_id(2);
  if (DST_X >= shared_int4_0.x || DST_Y >= 6 || DST_Z >= shared_int4_0.y) {
    return; 
  }
  int tile_x = (DST_X % shared_int4_0.z) * 4;
  int tile_y = (DST_X / shared_int4_0.z) * 4;
  ACCUM_FLT4 I0, I1, I2, I3, I4, I5;
  ACCUM_FLT bt_ar[6];
  ACCUM_FLT4 t0 = TO_ACCUM_TYPE(read_imagef(bt_tex2d, smp_none, (int2)(DST_Y * 2 + 0, 0)));
  ACCUM_FLT4 t1 = TO_ACCUM_TYPE(read_imagef(bt_tex2d, smp_none, (int2)(DST_Y * 2 + 1, 0)));
  DST_Y *= 6;
  bt_ar[0] = t0.x;
  bt_ar[1] = t0.y;
  bt_ar[2] = t0.z;
  bt_ar[3] = t0.w;
  bt_ar[4] = t1.x;
  bt_ar[5] = t1.y;
  int xc0 = tile_x + shared_int4_0.w + 0;
  ACCUM_FLT m0_x = (ACCUM_FLT)(xc0 >= 0 && xc0 < shared_int4_1.x);
  bool inx0 = (xc0 >= 0 && xc0 < shared_int4_1.x);
  xc0 = clamp(xc0, 0, shared_int4_1.x - 1);
  int src_a_0 = (((DST_Z) * shared_int4_1.y + (0)) * shared_int4_1.x + (xc0));;
  int xc1 = tile_x + shared_int4_0.w + 1;
  ACCUM_FLT m1_x = (ACCUM_FLT)(xc1 >= 0 && xc1 < shared_int4_1.x);
  bool inx1 = (xc1 >= 0 && xc1 < shared_int4_1.x);
  xc1 = clamp(xc1, 0, shared_int4_1.x - 1);
  int src_a_1 = (((DST_Z) * shared_int4_1.y + (0)) * shared_int4_1.x + (xc1));;
  int xc2 = tile_x + shared_int4_0.w + 2;
  ACCUM_FLT m2_x = (ACCUM_FLT)(xc2 >= 0 && xc2 < shared_int4_1.x);
  bool inx2 = (xc2 >= 0 && xc2 < shared_int4_1.x);
  xc2 = clamp(xc2, 0, shared_int4_1.x - 1);
  int src_a_2 = (((DST_Z) * shared_int4_1.y + (0)) * shared_int4_1.x + (xc2));;
  int xc3 = tile_x + shared_int4_0.w + 3;
  ACCUM_FLT m3_x = (ACCUM_FLT)(xc3 >= 0 && xc3 < shared_int4_1.x);
  bool inx3 = (xc3 >= 0 && xc3 < shared_int4_1.x);
  xc3 = clamp(xc3, 0, shared_int4_1.x - 1);
  int src_a_3 = (((DST_Z) * shared_int4_1.y + (0)) * shared_int4_1.x + (xc3));;
  int xc4 = tile_x + shared_int4_0.w + 4;
  ACCUM_FLT m4_x = (ACCUM_FLT)(xc4 >= 0 && xc4 < shared_int4_1.x);
  bool inx4 = (xc4 >= 0 && xc4 < shared_int4_1.x);
  xc4 = clamp(xc4, 0, shared_int4_1.x - 1);
  int src_a_4 = (((DST_Z) * shared_int4_1.y + (0)) * shared_int4_1.x + (xc4));;
  int xc5 = tile_x + shared_int4_0.w + 5;
  ACCUM_FLT m5_x = (ACCUM_FLT)(xc5 >= 0 && xc5 < shared_int4_1.x);
  bool inx5 = (xc5 >= 0 && xc5 < shared_int4_1.x);
  xc5 = clamp(xc5, 0, shared_int4_1.x - 1);
  int src_a_5 = (((DST_Z) * shared_int4_1.y + (0)) * shared_int4_1.x + (xc5));;
  {
    int yc = tile_y + shared_int4_1.z;
    bool iny = (yc >= 0 && yc < shared_int4_1.y);
    int offset = select(0, yc * shared_int4_1.x, iny);
    ACCUM_FLT bt = bt_ar[0] * (ACCUM_FLT)(iny);
    ACCUM_FLT4 src0 = src_tensor_buffer[src_a_0 + offset] * m0_x;
    I0 = bt * src0;
    ACCUM_FLT4 src1 = src_tensor_buffer[src_a_1 + offset] * m1_x;
    I1 = bt * src1;
    ACCUM_FLT4 src2 = src_tensor_buffer[src_a_2 + offset] * m2_x;
    I2 = bt * src2;
    ACCUM_FLT4 src3 = src_tensor_buffer[src_a_3 + offset] * m3_x;
    I3 = bt * src3;
    ACCUM_FLT4 src4 = src_tensor_buffer[src_a_4 + offset] * m4_x;
    I4 = bt * src4;
    ACCUM_FLT4 src5 = src_tensor_buffer[src_a_5 + offset] * m5_x;
    I5 = bt * src5;
  }
  {
    int yc = tile_y + shared_int4_1.z + (1);
    bool iny = (yc >= 0 && yc < shared_int4_1.y);
    int offset = select(0, yc * shared_int4_1.x, iny);
    ACCUM_FLT bt = bt_ar[1] * (ACCUM_FLT)(iny);
    ACCUM_FLT4 src0 = src_tensor_buffer[src_a_0 + offset] * m0_x;
    I0 += bt * src0;
    ACCUM_FLT4 src1 = src_tensor_buffer[src_a_1 + offset] * m1_x;
    I1 += bt * src1;
    ACCUM_FLT4 src2 = src_tensor_buffer[src_a_2 + offset] * m2_x;
    I2 += bt * src2;
    ACCUM_FLT4 src3 = src_tensor_buffer[src_a_3 + offset] * m3_x;
    I3 += bt * src3;
    ACCUM_FLT4 src4 = src_tensor_buffer[src_a_4 + offset] * m4_x;
    I4 += bt * src4;
    ACCUM_FLT4 src5 = src_tensor_buffer[src_a_5 + offset] * m5_x;
    I5 += bt * src5;
  }
  {
    int yc = tile_y + shared_int4_1.z + (2);
    bool iny = (yc >= 0 && yc < shared_int4_1.y);
    int offset = select(0, yc * shared_int4_1.x, iny);
    ACCUM_FLT bt = bt_ar[2] * (ACCUM_FLT)(iny);
    ACCUM_FLT4 src0 = src_tensor_buffer[src_a_0 + offset] * m0_x;
    I0 += bt * src0;
    ACCUM_FLT4 src1 = src_tensor_buffer[src_a_1 + offset] * m1_x;
    I1 += bt * src1;
    ACCUM_FLT4 src2 = src_tensor_buffer[src_a_2 + offset] * m2_x;
    I2 += bt * src2;
    ACCUM_FLT4 src3 = src_tensor_buffer[src_a_3 + offset] * m3_x;
    I3 += bt * src3;
    ACCUM_FLT4 src4 = src_tensor_buffer[src_a_4 + offset] * m4_x;
    I4 += bt * src4;
    ACCUM_FLT4 src5 = src_tensor_buffer[src_a_5 + offset] * m5_x;
    I5 += bt * src5;
  }
  {
    int yc = tile_y + shared_int4_1.z + (3);
    bool iny = (yc >= 0 && yc < shared_int4_1.y);
    int offset = select(0, yc * shared_int4_1.x, iny);
    ACCUM_FLT bt = bt_ar[3] * (ACCUM_FLT)(iny);
    ACCUM_FLT4 src0 = src_tensor_buffer[src_a_0 + offset] * m0_x;
    I0 += bt * src0;
    ACCUM_FLT4 src1 = src_tensor_buffer[src_a_1 + offset] * m1_x;
    I1 += bt * src1;
    ACCUM_FLT4 src2 = src_tensor_buffer[src_a_2 + offset] * m2_x;
    I2 += bt * src2;
    ACCUM_FLT4 src3 = src_tensor_buffer[src_a_3 + offset] * m3_x;
    I3 += bt * src3;
    ACCUM_FLT4 src4 = src_tensor_buffer[src_a_4 + offset] * m4_x;
    I4 += bt * src4;
    ACCUM_FLT4 src5 = src_tensor_buffer[src_a_5 + offset] * m5_x;
    I5 += bt * src5;
  }
  {
    int yc = tile_y + shared_int4_1.z + (4);
    bool iny = (yc >= 0 && yc < shared_int4_1.y);
    int offset = select(0, yc * shared_int4_1.x, iny);
    ACCUM_FLT bt = bt_ar[4] * (ACCUM_FLT)(iny);
    ACCUM_FLT4 src0 = src_tensor_buffer[src_a_0 + offset] * m0_x;
    I0 += bt * src0;
    ACCUM_FLT4 src1 = src_tensor_buffer[src_a_1 + offset] * m1_x;
    I1 += bt * src1;
    ACCUM_FLT4 src2 = src_tensor_buffer[src_a_2 + offset] * m2_x;
    I2 += bt * src2;
    ACCUM_FLT4 src3 = src_tensor_buffer[src_a_3 + offset] * m3_x;
    I3 += bt * src3;
    ACCUM_FLT4 src4 = src_tensor_buffer[src_a_4 + offset] * m4_x;
    I4 += bt * src4;
    ACCUM_FLT4 src5 = src_tensor_buffer[src_a_5 + offset] * m5_x;
    I5 += bt * src5;
  }
  {
    int yc = tile_y + shared_int4_1.z + (5);
    bool iny = (yc >= 0 && yc < shared_int4_1.y);
    int offset = select(0, yc * shared_int4_1.x, iny);
    ACCUM_FLT bt = bt_ar[5] * (ACCUM_FLT)(iny);
    ACCUM_FLT4 src0 = src_tensor_buffer[src_a_0 + offset] * m0_x;
    I0 += bt * src0;
    ACCUM_FLT4 src1 = src_tensor_buffer[src_a_1 + offset] * m1_x;
    I1 += bt * src1;
    ACCUM_FLT4 src2 = src_tensor_buffer[src_a_2 + offset] * m2_x;
    I2 += bt * src2;
    ACCUM_FLT4 src3 = src_tensor_buffer[src_a_3 + offset] * m3_x;
    I3 += bt * src3;
    ACCUM_FLT4 src4 = src_tensor_buffer[src_a_4 + offset] * m4_x;
    I4 += bt * src4;
    ACCUM_FLT4 src5 = src_tensor_buffer[src_a_5 + offset] * m5_x;
    I5 += bt * src5;
  }

  // if (DST_X == 893 && DST_Y/6 == 2 && DST_Z == 15) {
  //   printf("TFLITE WINOGRAD\n");
  //   printf("%f, %f, %f, %f\n", I0.x, I0.y, I0.z, I0.w);
  //   printf("%f, %f, %f, %f\n", I1.x, I1.y, I1.z, I1.w);
  //   printf("%f, %f, %f, %f\n", I2.x, I2.y, I2.z, I2.w);
  //   printf("%f, %f, %f, %f\n", I3.x, I3.y, I3.z, I3.w);
  //   printf("%f, %f, %f, %f\n", I4.x, I4.y, I4.z, I4.w);
  //   printf("%f, %f, %f, %f\n", I5.x, I5.y, I5.z, I5.w);
  // }

  {
    FLT4 r0 = TO_FLT4(I0 + Bt[2] * I2 + Bt[4] * I4);
    dst_tensor_buffer[(((DST_Z) * shared_int4_1.w + (DST_Y)) * shared_int4_2.x + (DST_X))] = r0;
;
    DST_Y++;
  }
  {
    FLT4 r0 = TO_FLT4(Bt[7] * I1 + Bt[8] * I2 + Bt[9] * I3 + Bt[10] * I4);
    dst_tensor_buffer[(((DST_Z) * shared_int4_1.w + (DST_Y)) * shared_int4_2.x + (DST_X))] = r0;
;
    DST_Y++;
  }
  {
    FLT4 r0 = TO_FLT4(Bt[13] * I1 + Bt[14] * I2 + Bt[15] * I3 + Bt[16] * I4);
    dst_tensor_buffer[(((DST_Z) * shared_int4_1.w + (DST_Y)) * shared_int4_2.x + (DST_X))] = r0;
;
    DST_Y++;
  }
  {
    FLT4 r0 = TO_FLT4(Bt[19] * I1 + Bt[20] * I2 + Bt[21] * I3 + Bt[22] * I4);
    dst_tensor_buffer[(((DST_Z) * shared_int4_1.w + (DST_Y)) * shared_int4_2.x + (DST_X))] = r0;
;
    DST_Y++;
  }
  {
    FLT4 r0 = TO_FLT4(Bt[25] * I1 + Bt[26] * I2 + Bt[27] * I3 + Bt[28] * I4);
    dst_tensor_buffer[(((DST_Z) * shared_int4_1.w + (DST_Y)) * shared_int4_2.x + (DST_X))] = r0;
;
    DST_Y++;
  }
  {
    FLT4 r0 = TO_FLT4(Bt[31] * I1 + Bt[33] * I3 + I5);
    dst_tensor_buffer[(((DST_Z) * shared_int4_1.w + (DST_Y)) * shared_int4_2.x + (DST_X))] = r0;
;
    DST_Y++;
  }
}